Skip to content

[OpenACC][CIR] 'update' construct lowering + a few clauses #146378

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom
Merged
merged 2 commits into from
Jun 30, 2025
Merged

[OpenACC][CIR] 'update' construct lowering + a few clauses #146378

merged 2 commits into from
Jun 30, 2025

Conversation

@erichkeane
Copy link
Collaborator

@erichkeane erichkeane commented Jun 30, 2025

The 'update' construct has 3 'var-list' clauses, device, self, and host. Each has a pretty simple data-operand type syntax in the IR, so this patch implements them as well. At least one of those is required to be present on an 'update', so we cannot do any lowering without them.

Note that 'self' and 'host' are aliases.

@erichkeane
[OpenACC][CIR] 'update' construct lowering + a few clauses
87565b6 
The 'update' construct has 3 'var-list' clauses, device, self, and host.
Each has a pretty simple data-operand type syntax in the IR, so this
patch implements them as well.  At least one of those is required to be
present on an 'update', so we cannot do any lowering without them.

Note that 'self' and 'host' are aliases.
@llvmbot llvmbot added clang Clang issues not falling into any other category ClangIR Anything related to the ClangIR project labels Jun 30, 2025
@llvmbot
Copy link
Member

llvmbot commented Jun 30, 2025 •
edited
Loading

@llvm/pr-subscribers-clangir

@llvm/pr-subscribers-clang

Author: Erich Keane (erichkeane)

Changes

The 'update' construct has 3 'var-list' clauses, device, self, and host. Each has a pretty simple data-operand type syntax in the IR, so this patch implements them as well. At least one of those is required to be present on an 'update', so we cannot do any lowering without them.

Note that 'self' and 'host' are aliases.

Full diff: https://github.com/llvm/llvm-project/pull/146378.diff

3 Files Affected:

  • (modified) clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp (+30-3)
  • (modified) clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp (+4-2)
  • (added) clang/test/CIR/CodeGenOpenACC/update.c (+67)
diff --git a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
index cc0f3b77c1a65..b7a73e2f62945 100644
--- a/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenOpenACCClause.cpp
@@ -613,12 +613,39 @@ class OpenACCClauseCIREmitter final
       } else {
         llvm_unreachable("var-list version of self shouldn't get here");
       }
+    } else if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
+      assert(!clause.isEmptySelfClause() && !clause.isConditionExprClause() &&
+             "var-list version of self required for update");
+      for (const Expr *var : clause.getVarList())
+        addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
+            var, mlir::acc::DataClause::acc_update_self, {},
+            /*structured=*/false, /*implicit=*/false);
     } else if constexpr (isCombinedType<OpTy>) {
       applyToComputeOp(clause);
     } else {
-      // TODO: When we've implemented this for everything, switch this to an
-      // unreachable. update construct remains.
-      return clauseNotImplemented(clause);
+      llvm_unreachable("Unknown construct kind in VisitSelfClause");
+    }
+  }
+
+  void VisitHostClause(const OpenACCHostClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
+      for (const Expr *var : clause.getVarList())
+        addDataOperand<mlir::acc::GetDevicePtrOp, mlir::acc::UpdateHostOp>(
+            var, mlir::acc::DataClause::acc_update_host, {},
+            /*structured=*/false, /*implicit=*/false);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitHostClause");
+    }
+  }
+
+  void VisitDeviceClause(const OpenACCDeviceClause &clause) {
+    if constexpr (isOneOfTypes<OpTy, mlir::acc::UpdateOp>) {
+      for (const Expr *var : clause.getVarList())
+        addDataOperand<mlir::acc::UpdateDeviceOp>(
+            var, mlir::acc::DataClause::acc_update_device, {},
+            /*structured=*/false, /*implicit=*/false);
+    } else {
+      llvm_unreachable("Unknown construct kind in VisitDeviceClause");
     }
   }
 
@@ -1095,6 +1122,7 @@ EXPL_SPEC(mlir::acc::WaitOp)
 EXPL_SPEC(mlir::acc::HostDataOp)
 EXPL_SPEC(mlir::acc::EnterDataOp)
 EXPL_SPEC(mlir::acc::ExitDataOp)
+EXPL_SPEC(mlir::acc::UpdateOp)
 #undef EXPL_SPEC
 
 template <typename ComputeOp, typename LoopOp>
diff --git a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
index f3a635b7c83eb..5993056bf06ba 100644
--- a/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenStmtOpenACC.cpp
@@ -266,8 +266,10 @@ mlir::LogicalResult CIRGenFunction::emitOpenACCExitDataConstruct(
 
 mlir::LogicalResult
 CIRGenFunction::emitOpenACCUpdateConstruct(const OpenACCUpdateConstruct &s) {
-  cgm.errorNYI(s.getSourceRange(), "OpenACC Update Construct");
-  return mlir::failure();
+  mlir::Location start = getLoc(s.getSourceRange().getBegin());
+  emitOpenACCOp<UpdateOp>(start, s.getDirectiveKind(), s.getDirectiveLoc(),
+                          s.clauses());
+  return mlir::success();
 }
 mlir::LogicalResult
 CIRGenFunction::emitOpenACCAtomicConstruct(const OpenACCAtomicConstruct &s) {
diff --git a/clang/test/CIR/CodeGenOpenACC/update.c b/clang/test/CIR/CodeGenOpenACC/update.c
new file mode 100644
index 0000000000000..4e25a1df2a42b
--- /dev/null
+++ b/clang/test/CIR/CodeGenOpenACC/update.c
@@ -0,0 +1,67 @@
+// RUN: %clang_cc1 -fopenacc -Wno-openacc-self-if-potential-conflict -emit-cir -fclangir %s -o - | FileCheck %s
+
+void acc_update(int parmVar, int *ptrParmVar) {
+  // CHECK: cir.func{{.*}} @acc_update(%[[ARG:.*]]: !s32i{{.*}}, %[[PTRARG:.*]]: !cir.ptr<!s32i>{{.*}}) {
+  // CHECK-NEXT: %[[PARM:.*]] = cir.alloca !s32i, !cir.ptr<!s32i>, ["parmVar", init]
+  // CHECK-NEXT: %[[PTRPARM:.*]] = cir.alloca !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>, ["ptrParmVar", init]
+  // CHECK-NEXT: cir.store %[[ARG]], %[[PARM]] : !s32i, !cir.ptr<!s32i>
+  // CHECK-NEXT: cir.store %[[PTRARG]], %[[PTRPARM]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>
+
+#pragma acc update device(parmVar)
+  // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]] : !cir.ptr<!s32i>)
+
+#pragma acc update device(parmVar, ptrParmVar)
+  // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+
+#pragma acc update device(parmVar) device(ptrParmVar)
+  // CHECK-NEXT: %[[UPD_DEV1:.*]] = acc.update_device varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[UPD_DEV1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+
+#pragma acc update host(parmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_host>, name = "parmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc update host(parmVar, ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_host>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_host>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc update host(parmVar) host(ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_host>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_host>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {name = "parmVar", structured = false}
+
+#pragma acc update self(parmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]] : !cir.ptr<!s32i>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
+
+#pragma acc update self(parmVar, ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
+
+#pragma acc update self(parmVar) self(ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[GDP2:.*]] = acc.getdeviceptr varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[GDP2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP2]] : !cir.ptr<!cir.ptr<!s32i>>) to varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) {dataClause = #acc<data_clause acc_update_self>, name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
+
+#pragma acc update self(parmVar) device(ptrParmVar)
+  // CHECK-NEXT: %[[GDP1:.*]] = acc.getdeviceptr varPtr(%[[PARM]] : !cir.ptr<!s32i>) -> !cir.ptr<!s32i> {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
+  // CHECK-NEXT: %[[UPD_DEV2:.*]] = acc.update_device varPtr(%[[PTRPARM]] : !cir.ptr<!cir.ptr<!s32i>>) -> !cir.ptr<!cir.ptr<!s32i>> {name = "ptrParmVar", structured = false}
+  // CHECK-NEXT: acc.update dataOperands(%[[GDP1]], %[[UPD_DEV2]] : !cir.ptr<!s32i>, !cir.ptr<!cir.ptr<!s32i>>)
+  // CHECK-NEXT: acc.update_host accPtr(%[[GDP1]] : !cir.ptr<!s32i>) to varPtr(%[[PARM]] : !cir.ptr<!s32i>) {dataClause = #acc<data_clause acc_update_self>, name = "parmVar", structured = false}
+}

clementval
clementval approved these changes Jun 30, 2025
View reviewed changes
Copy link
Contributor

@clementval clementval left a comment

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

LGTM

@erichkeane erichkeane merged commit 125dbe1 into llvm:main Jun 30, 2025
7 checks passed
@llvm-ci
Copy link
Collaborator

llvm-ci commented Jul 1, 2025

LLVM Buildbot has detected a new failure on builder llvm-x86_64-debian-dylib running on gribozavr4 while building clang at step 7 "test-build-unified-tree-check-llvm".

Full details are available at: https://lab.llvm.org/buildbot/#/builders/60/builds/31714

Here is the relevant piece of the build log for the reference 
Step 7 (test-build-unified-tree-check-llvm) failure: test (failure)
******************** TEST 'LLVM :: tools/llvm-exegesis/RISCV/rvv/filter.test' FAILED ********************
Exit Code: 2

Command Output (stderr):
--
/b/1/llvm-x86_64-debian-dylib/build/bin/llvm-exegesis -mtriple=riscv64 -mcpu=sifive-x280 -benchmark-phase=assemble-measured-code --mode=inverse_throughput --opcode-name=PseudoVNCLIPU_WX_M1_MASK     --riscv-filter-config='vtype = {VXRM: rod, AVL: VLMAX, SEW: e(8|16), Policy: ta/mu}' --max-configs-per-opcode=1000 --min-instructions=10 | /b/1/llvm-x86_64-debian-dylib/build/bin/FileCheck /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/tools/llvm-exegesis/RISCV/rvv/filter.test # RUN: at line 1
+ /b/1/llvm-x86_64-debian-dylib/build/bin/llvm-exegesis -mtriple=riscv64 -mcpu=sifive-x280 -benchmark-phase=assemble-measured-code --mode=inverse_throughput --opcode-name=PseudoVNCLIPU_WX_M1_MASK '--riscv-filter-config=vtype = {VXRM: rod, AVL: VLMAX, SEW: e(8|16), Policy: ta/mu}' --max-configs-per-opcode=1000 --min-instructions=10
PseudoVNCLIPU_WX_M1_MASK: Failed to produce any snippet via: instruction has tied variables, avoiding Read-After-Write issue, picking random def and use registers not aliasing each other, for uses, one unique register for each position
+ /b/1/llvm-x86_64-debian-dylib/build/bin/FileCheck /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/tools/llvm-exegesis/RISCV/rvv/filter.test
FileCheck error: '<stdin>' is empty.
FileCheck command line:  /b/1/llvm-x86_64-debian-dylib/build/bin/FileCheck /b/1/llvm-x86_64-debian-dylib/llvm-project/llvm/test/tools/llvm-exegesis/RISCV/rvv/filter.test

--

********************

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Reviewers

@clementval clementval clementval approved these changes

@razvanlupusoru razvanlupusoru Awaiting requested review from razvanlupusoru

@lanza lanza Awaiting requested review from lanza lanza is a code owner

@bcardosolopes bcardosolopes Awaiting requested review from bcardosolopes bcardosolopes is a code owner

Assignees

No one assigned

Labels

clang Clang issues not falling into any other category ClangIR Anything related to the ClangIR project

Projects

None yet

Milestone

No milestone

Development

Successfully merging this pull request may close these issues.

4 participants

@erichkeane @llvmbot @llvm-ci @clementval